#define K_MAX_SHAPE_DIM 0
#include "kernel_operator.h"
using namespace AscendC;
constexpr int32_t BUFFER_NUM = 2;
class KernelGluGradJvp_float {
public:
    __aicore__ inline KernelGluGradJvp_float() {}
    __aicore__ inline void Init(GM_ADDR x_grad, GM_ADDR y_grad, GM_ADDR x, GM_ADDR v_y, GM_ADDR v_x, GM_ADDR jvp_out,
                                uint32_t bigNum, uint32_t bigLength, uint32_t smallLength, 
                                uint32_t interval, TPipe* pipeIn) {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
        uint32_t coreDataNum;
        if (GetBlockIdx() < bigNum) {
            coreDataNum = bigLength;
            xGradGm.SetGlobalBuffer((__gm__ float*)x_grad + bigLength * interval * 2 * GetBlockIdx());
            vyGm.SetGlobalBuffer   ((__gm__ float*)v_y    + bigLength * interval * GetBlockIdx());
            vxGm.SetGlobalBuffer   ((__gm__ float*)v_x    + bigLength * interval * 2 * GetBlockIdx());
            outGm.SetGlobalBuffer ((__gm__ float*)jvp_out + bigLength * interval * 2 * GetBlockIdx());
        } else {
            coreDataNum = smallLength;
            xGradGm.SetGlobalBuffer((__gm__ float*)x_grad + bigLength * interval * 2 * bigNum + smallLength * interval * 2 * (GetBlockIdx() - bigNum));
            vyGm.SetGlobalBuffer   ((__gm__ float*)v_y    + bigLength * interval * bigNum + smallLength * interval * (GetBlockIdx() - bigNum));
            vxGm.SetGlobalBuffer   ((__gm__ float*)v_x    + bigLength * interval * 2 * bigNum + smallLength * interval * 2 * (GetBlockIdx() - bigNum));
            outGm.SetGlobalBuffer ((__gm__ float*)jvp_out + bigLength * interval * 2 * bigNum + smallLength * interval * 2 * (GetBlockIdx() - bigNum));
        }
        this->repeat = coreDataNum;
        this->interval = interval;
        this->tileDataLenght = 3488;
        this->loopCount = this->interval/this->tileDataLenght;
        this->tailDataLenght = this->interval - this->tileDataLenght*this->loopCount;
        if(this->tailDataLenght == 0){this->tailDataLenght = this->tileDataLenght; this->loopCount--;}
        this->tailDataLenght_32 = (this->tailDataLenght * sizeof(float) + 31)/32*32 /sizeof(float);
        pipeIn->InitBuffer(inQueueXGrad, BUFFER_NUM, this->tileDataLenght *2 * sizeof(float));
        pipeIn->InitBuffer(inQueueVy, BUFFER_NUM, this->tileDataLenght * sizeof(float));
        pipeIn->InitBuffer(inQueueVx, BUFFER_NUM, this->tileDataLenght *2 * sizeof(float));
        pipeIn->InitBuffer(outQueueOut, BUFFER_NUM, this->tileDataLenght *2 * sizeof(float));
    }
    __aicore__ inline void Process() {
        for (int32_t i = 0; i < this->repeat; i++) {
            this->processDataLenght = this->tileDataLenght;
            for(int32_t j=0; j<this->loopCount; j++)
            {
                CopyIn(i, j);
                Compute(i);
                CopyOut(i, j);
            }
            this->processDataLenght = this->tailDataLenght;
            CopyInTail(i, this->loopCount);
            Compute(i);
            CopyOut(i, this->loopCount);
        }
    }
private:
        __aicore__ inline void CopyIn(int32_t progress, int32_t step) {
        LocalTensor<float> xGradLocal = inQueueXGrad.AllocTensor<float>();
        LocalTensor<float> vyLocal = inQueueVy.AllocTensor<float>();
        LocalTensor<float> vxLocal = inQueueVx.AllocTensor<float>();
        DataCopy(xGradLocal, xGradGm[progress * this->interval * 2 + step*this->tileDataLenght], this->processDataLenght);
        DataCopy(xGradLocal[this->tileDataLenght], xGradGm[progress * this->interval * 2 + this->interval + step*this->tileDataLenght], this->processDataLenght);
        DataCopy(vyLocal, vyGm[progress * this->interval + step*this->tileDataLenght], this->processDataLenght);
        DataCopy(vxLocal, vxGm[progress * this->interval * 2 + step*this->tileDataLenght], this->processDataLenght);
        DataCopy(vxLocal[this->tileDataLenght], vxGm[progress * this->interval * 2 + this->interval + step*this->tileDataLenght], this->processDataLenght);
        inQueueXGrad.EnQue(xGradLocal);
        inQueueVy.EnQue(vyLocal);
        inQueueVx.EnQue(vxLocal);
    }
    __aicore__ inline void CopyInTail(int32_t progress, int32_t step) {
        LocalTensor<float> xGradLocal = inQueueXGrad.AllocTensor<float>();
        LocalTensor<float> vyLocal = inQueueVy.AllocTensor<float>();
        LocalTensor<float> vxLocal = inQueueVx.AllocTensor<float>();
        DataCopy(xGradLocal, xGradGm[progress * this->interval * 2 + step*this->tileDataLenght], this->tailDataLenght_32);
        DataCopy(xGradLocal[this->tileDataLenght], xGradGm[progress * this->interval * 2 + this->interval + step*this->tileDataLenght], this->tailDataLenght_32);
        DataCopy(vyLocal, vyGm[progress * this->interval + step*this->tileDataLenght], this->tailDataLenght_32);
        DataCopy(vxLocal, vxGm[progress * this->interval * 2 + step*this->tileDataLenght], this->tailDataLenght_32);
        DataCopy(vxLocal[this->tileDataLenght], vxGm[progress * this->interval * 2 + this->interval + step*this->tileDataLenght], this->tailDataLenght_32);
        inQueueXGrad.EnQue(xGradLocal);
        inQueueVy.EnQue(vyLocal);
        inQueueVx.EnQue(vxLocal);
    }   
    __aicore__ inline void Compute(int32_t progress) {
        LocalTensor<float> xGradLocal = inQueueXGrad.DeQue<float>();
        LocalTensor<float> vyLocal = inQueueVy.DeQue<float>();
        LocalTensor<float> vxLocal = inQueueVx.DeQue<float>();
        LocalTensor<float> outLocal = outQueueOut.AllocTensor<float>();
        Mul(outLocal, vxLocal[this->tileDataLenght], xGradLocal, this->processDataLenght);
        Sub(outLocal, vxLocal[this->tileDataLenght], outLocal, this->processDataLenght);
        Add(outLocal[this->tileDataLenght], outLocal, vyLocal, this->processDataLenght);
        Mul(outLocal, outLocal[this->tileDataLenght], xGradLocal, this->processDataLenght);
        Mul(outLocal[this->tileDataLenght], outLocal[this->tileDataLenght], xGradLocal[this->tileDataLenght], this->processDataLenght);
        Mul(vxLocal[this->tileDataLenght], vxLocal[this->tileDataLenght], xGradLocal[this->tileDataLenght], this->processDataLenght);
        Mul(xGradLocal[this->tileDataLenght], vxLocal, xGradLocal, this->processDataLenght);
        Sub(vxLocal, vxLocal, xGradLocal[this->tileDataLenght], this->processDataLenght);
        Sub(vxLocal, vxLocal, vxLocal[this->tileDataLenght], this->processDataLenght);
        Mul(vxLocal, vxLocal, xGradLocal, this->processDataLenght);
        Add(outLocal[this->tileDataLenght], outLocal[this->tileDataLenght], vxLocal, this->processDataLenght);
        inQueueXGrad.FreeTensor(xGradLocal);
        inQueueVy.FreeTensor(vyLocal);
        inQueueVx.FreeTensor(vxLocal);
        outQueueOut.EnQue<float>(outLocal);
    }
    __aicore__ inline void CopyOut(int32_t progress, int32_t step) {
        LocalTensor<float> outLocal = outQueueOut.DeQue<float>();
        DataCopyExtParams copyParams{1, static_cast<uint32_t>(this->processDataLenght * sizeof(float)), 0, 0, 0}; 
        DataCopyPad(outGm[progress * this->interval * 2 + step*this->tileDataLenght], outLocal, copyParams);
        DataCopyPad(outGm[progress * this->interval * 2 + this->interval + step*this->tileDataLenght], outLocal[this->tileDataLenght], copyParams);
        outQueueOut.FreeTensor(outLocal);
    }
private:
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueXGrad, inQueueVy, inQueueVx;
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueOut;
    GlobalTensor<float> xGradGm;
    GlobalTensor<float> vyGm;
    GlobalTensor<float> vxGm;
    GlobalTensor<float> outGm;
    uint32_t repeat;
    uint32_t interval;
    uint32_t loopCount;
    uint32_t tileDataLenght;
    uint32_t tailDataLenght;
    uint32_t tailDataLenght_32;
    uint32_t processDataLenght;
};
class KernelGluGradJvp_float1 {
public:
    __aicore__ inline KernelGluGradJvp_float1() {}
    __aicore__ inline void Init(GM_ADDR x_grad, GM_ADDR y_grad, GM_ADDR x, GM_ADDR v_y, GM_ADDR v_x, GM_ADDR jvp_out,
                                uint32_t bigNum, uint32_t bigLength, uint32_t smallLength, 
                                uint32_t interval, TPipe* pipeIn) {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
        uint32_t coreDataNum;
        if (GetBlockIdx() < bigNum) {
            coreDataNum = bigLength;
            xGradGm.SetGlobalBuffer((__gm__ float*)x_grad + bigLength * interval * 2 * GetBlockIdx());
            vyGm.SetGlobalBuffer   ((__gm__ float*)v_y    + bigLength * interval * GetBlockIdx());
            vxGm.SetGlobalBuffer   ((__gm__ float*)v_x    + bigLength * interval * 2 * GetBlockIdx());
            outGm.SetGlobalBuffer ((__gm__ float*)jvp_out + bigLength * interval * 2 * GetBlockIdx());
        } else {
            coreDataNum = smallLength;
            xGradGm.SetGlobalBuffer((__gm__ float*)x_grad + bigLength * interval * 2 * bigNum + smallLength * interval * 2 * (GetBlockIdx() - bigNum));
            vyGm.SetGlobalBuffer   ((__gm__ float*)v_y    + bigLength * interval * bigNum + smallLength * interval * (GetBlockIdx() - bigNum));
            vxGm.SetGlobalBuffer   ((__gm__ float*)v_x    + bigLength * interval * 2 * bigNum + smallLength * interval * 2 * (GetBlockIdx() - bigNum));
            outGm.SetGlobalBuffer ((__gm__ float*)jvp_out + bigLength * interval * 2 * bigNum + smallLength * interval * 2 * (GetBlockIdx() - bigNum));
        }
        this->interval_32 = (interval * sizeof(float) + 31)/32*32 /sizeof(float);
        this->repeat = coreDataNum;
        this->interval = interval;
        this->tileDataLenght = 3488/this->interval_32;
        this->loopCount = this->repeat/this->tileDataLenght;
        this->tailDataLenght = this->repeat - this->tileDataLenght*this->loopCount;
        if(this->tailDataLenght == 0){this->tailDataLenght = this->tileDataLenght; this->loopCount--;}
        pipeIn->InitBuffer(inQueueXGrad, BUFFER_NUM, 3488 *2 * sizeof(float));
        pipeIn->InitBuffer(inQueueVy, BUFFER_NUM, 3488 * sizeof(float));
        pipeIn->InitBuffer(inQueueVx, BUFFER_NUM, 3488 *2 * sizeof(float));
        pipeIn->InitBuffer(outQueueOut, BUFFER_NUM, 3488 *2 * sizeof(float));
    }
    __aicore__ inline void Process() {
        this->processDataLenght = this->tileDataLenght * this->interval_32;
        for (int32_t i = 0; i < this->loopCount; i++) {
            CopyIn(i);
            Compute(i);
            CopyOut(i);
        }
        this->processDataLenght = this->tailDataLenght * this->interval_32;
        CopyInTail(this->loopCount);
        Compute(this->loopCount);
        CopyOutTail(this->loopCount);
    }
private:
    __aicore__ inline void CopyIn(int32_t progress) {
        DataCopyExtParams copyParams{static_cast<uint16_t>(this->tileDataLenght), static_cast<uint32_t>(this->interval * sizeof(float)), 0, 0, 0};
        DataCopyExtParams copyParams1{static_cast<uint16_t>(this->tileDataLenght), static_cast<uint32_t>(this->interval * sizeof(float)), static_cast<uint32_t>(this->interval * sizeof(float)), 0, 0};
        DataCopyPadExtParams<float> padParams{false, 0, 0, 0};
        LocalTensor<float> xGradLocal = inQueueXGrad.AllocTensor<float>();
        LocalTensor<float> vyLocal = inQueueVy.AllocTensor<float>();
        LocalTensor<float> vxLocal = inQueueVx.AllocTensor<float>();
        DataCopyPad(xGradLocal, xGradGm[progress * this->interval * this->tileDataLenght * 2], copyParams1, padParams);
        DataCopyPad(xGradLocal[3488], xGradGm[progress * this->interval * this->tileDataLenght * 2 + this->interval], copyParams1, padParams);
        DataCopyPad(vyLocal, vyGm[progress * this->interval * this->tileDataLenght], copyParams, padParams);
        DataCopyPad(vxLocal, vxGm[progress * this->interval * this->tileDataLenght * 2], copyParams1, padParams);
        DataCopyPad(vxLocal[3488], vxGm[progress * this->interval * this->tileDataLenght * 2 + this->interval], copyParams1, padParams);
        inQueueXGrad.EnQue(xGradLocal);
        inQueueVy.EnQue(vyLocal);
        inQueueVx.EnQue(vxLocal);
    }
    __aicore__ inline void CopyInTail(int32_t progress) {
        DataCopyExtParams copyParams{static_cast<uint16_t>(this->tailDataLenght), static_cast<uint32_t>(this->interval * sizeof(float)), 0, 0, 0};
        DataCopyExtParams copyParams1{static_cast<uint16_t>(this->tailDataLenght), static_cast<uint32_t>(this->interval * sizeof(float)), static_cast<uint32_t>(this->interval * sizeof(float)), 0, 0};
        DataCopyPadExtParams<float> padParams{false, 0, 0, 0};
        LocalTensor<float> xGradLocal = inQueueXGrad.AllocTensor<float>();
        LocalTensor<float> vyLocal = inQueueVy.AllocTensor<float>();
        LocalTensor<float> vxLocal = inQueueVx.AllocTensor<float>();
        DataCopyPad(xGradLocal, xGradGm[progress * this->interval * this->tileDataLenght * 2], copyParams1, padParams);
        DataCopyPad(xGradLocal[3488], xGradGm[progress * this->interval * this->tileDataLenght * 2 + this->interval], copyParams1, padParams);
        DataCopyPad(vyLocal, vyGm[progress * this->interval * this->tileDataLenght], copyParams, padParams);
        DataCopyPad(vxLocal, vxGm[progress * this->interval * this->tileDataLenght * 2], copyParams1, padParams);
        DataCopyPad(vxLocal[3488], vxGm[progress * this->interval * this->tileDataLenght * 2 + this->interval], copyParams1, padParams);
        inQueueXGrad.EnQue(xGradLocal);
        inQueueVy.EnQue(vyLocal);
        inQueueVx.EnQue(vxLocal);
    }  
    __aicore__ inline void Compute(int32_t progress) {
        LocalTensor<float> xGradLocal = inQueueXGrad.DeQue<float>();
        LocalTensor<float> vyLocal = inQueueVy.DeQue<float>();
        LocalTensor<float> vxLocal = inQueueVx.DeQue<float>();
        LocalTensor<float> outLocal = outQueueOut.AllocTensor<float>();
        Mul(outLocal, vxLocal[3488], xGradLocal, this->processDataLenght);
        Sub(outLocal, vxLocal[3488], outLocal, this->processDataLenght);
        Add(outLocal[3488], outLocal, vyLocal, this->processDataLenght);
        Mul(outLocal, outLocal[3488], xGradLocal, this->processDataLenght);
        Mul(outLocal[3488], outLocal[3488], xGradLocal[3488], this->processDataLenght);
        Mul(vxLocal[3488], vxLocal[3488], xGradLocal[3488], this->processDataLenght);
        Mul(xGradLocal[3488], vxLocal, xGradLocal, this->processDataLenght);
        Sub(vxLocal, vxLocal, xGradLocal[3488], this->processDataLenght);
        Sub(vxLocal, vxLocal, vxLocal[3488], this->processDataLenght);
        Mul(vxLocal, vxLocal, xGradLocal, this->processDataLenght);
        Add(outLocal[3488], outLocal[3488], vxLocal, this->processDataLenght);
        inQueueXGrad.FreeTensor(xGradLocal);
        inQueueVy.FreeTensor(vyLocal);
        inQueueVx.FreeTensor(vxLocal);
        outQueueOut.EnQue<float>(outLocal);
    }
    __aicore__ inline void CopyOut(int32_t progress) {
        LocalTensor<float> outLocal = outQueueOut.DeQue<float>();
        DataCopyExtParams copyParams{static_cast<uint16_t>(this->tileDataLenght), static_cast<uint32_t>(this->interval * sizeof(float)), 0, static_cast<uint32_t>(this->interval * sizeof(float)), 0}; 
        DataCopyPad(outGm[progress * this->interval * this->tileDataLenght * 2], outLocal, copyParams);
        DataCopyPad(outGm[progress * this->interval * this->tileDataLenght * 2 + this->interval], outLocal[3488], copyParams);
        outQueueOut.FreeTensor(outLocal);
    }
    __aicore__ inline void CopyOutTail(int32_t progress) {
        LocalTensor<float> outLocal = outQueueOut.DeQue<float>();
        DataCopyExtParams copyParams{static_cast<uint16_t>(this->tailDataLenght), static_cast<uint32_t>(this->interval * sizeof(float)), 0, static_cast<uint32_t>(this->interval * sizeof(float)), 0}; 
        DataCopyPad(outGm[progress * this->interval * this->tileDataLenght * 2], outLocal, copyParams);
        DataCopyPad(outGm[progress * this->interval * this->tileDataLenght * 2 + this->interval], outLocal[3488], copyParams);
        outQueueOut.FreeTensor(outLocal);
    }
private:
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueXGrad, inQueueVy, inQueueVx;
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueOut;
    GlobalTensor<float> xGradGm;
    GlobalTensor<float> vyGm;
    GlobalTensor<float> vxGm;
    GlobalTensor<float> outGm;
    uint32_t repeat;
    uint32_t interval;
    uint32_t interval_32;
    uint32_t loopCount;
    uint32_t tileDataLenght;
    uint32_t tailDataLenght;
    uint32_t tailDataLenght_32;
    uint32_t processDataLenght;
};
class KernelGluGradJvp_half {
public:
    __aicore__ inline KernelGluGradJvp_half() {}
    __aicore__ inline void Init(GM_ADDR x_grad, GM_ADDR y_grad, GM_ADDR x, GM_ADDR v_y, GM_ADDR v_x, GM_ADDR jvp_out,
                                uint32_t bigNum, uint32_t bigLength, uint32_t smallLength, 
                                uint32_t interval, TPipe* pipeIn) {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
        uint32_t coreDataNum;
        if (GetBlockIdx() < bigNum) {
            coreDataNum = bigLength;
            xGradGm.SetGlobalBuffer((__gm__ DTYPE_X*)x_grad + bigLength * interval * 2 * GetBlockIdx());
            xGm.SetGlobalBuffer   ((__gm__ DTYPE_X*)x + bigLength * interval * 2 * GetBlockIdx());
            vyGm.SetGlobalBuffer  ((__gm__ DTYPE_X*)v_y    + bigLength * interval * GetBlockIdx());
            vxGm.SetGlobalBuffer  ((__gm__ DTYPE_X*)v_x    + bigLength * interval * 2 * GetBlockIdx());
            outGm.SetGlobalBuffer ((__gm__ DTYPE_X*)jvp_out + bigLength * interval * 2 * GetBlockIdx());
        } else {
            coreDataNum = smallLength;
            xGradGm.SetGlobalBuffer((__gm__ DTYPE_X*)x_grad + bigLength * interval * 2 * bigNum + smallLength * interval * 2 * (GetBlockIdx() - bigNum));
            xGm.SetGlobalBuffer   ((__gm__ DTYPE_X*)x + bigLength * interval * 2 * bigNum + smallLength * interval * 2 * (GetBlockIdx() - bigNum));
            vyGm.SetGlobalBuffer  ((__gm__ DTYPE_X*)v_y    + bigLength * interval * bigNum + smallLength * interval * (GetBlockIdx() - bigNum));
            vxGm.SetGlobalBuffer  ((__gm__ DTYPE_X*)v_x    + bigLength * interval * 2 * bigNum + smallLength * interval * 2 * (GetBlockIdx() - bigNum));
            outGm.SetGlobalBuffer ((__gm__ DTYPE_X*)jvp_out + bigLength * interval * 2 * bigNum + smallLength * interval * 2 * (GetBlockIdx() - bigNum));
        }
        this->repeat = coreDataNum;
        this->interval = interval;
        this->tileDataLenght = 4448;
        this->loopCount = this->interval/this->tileDataLenght;
        this->tailDataLenght = this->interval - this->tileDataLenght*this->loopCount;
        if(this->tailDataLenght == 0){this->tailDataLenght = this->tileDataLenght; this->loopCount--;}
        this->tailDataLenght_32 = (this->tailDataLenght * sizeof(DTYPE_X) + 31)/32*32 /sizeof(DTYPE_X);
        pipeIn->InitBuffer(inQueueX, BUFFER_NUM, this->tileDataLenght *2 * sizeof(DTYPE_X));
        pipeIn->InitBuffer(inQueueVy, BUFFER_NUM, this->tileDataLenght * 2 * sizeof(DTYPE_X));
        pipeIn->InitBuffer(inQueueVx, BUFFER_NUM, this->tileDataLenght *2 * sizeof(DTYPE_X));
        pipeIn->InitBuffer(outQueueOut, BUFFER_NUM, this->tileDataLenght *2 * sizeof(DTYPE_X));
        pipeIn->InitBuffer(QueueTmp1, this->tileDataLenght * sizeof(float));
        pipeIn->InitBuffer(QueueTmp2, this->tileDataLenght * sizeof(float));
        pipeIn->InitBuffer(QueueTmp3, this->tileDataLenght * sizeof(float));
    }
    __aicore__ inline void Process() {
        for (int32_t i = 0; i < this->repeat; i++) {
            this->processDataLenght = this->tileDataLenght;
            for(int32_t j=0; j<this->loopCount; j++)
            {
                CopyIn(i, j);
                Compute(i);
                CopyOut(i, j);
            }
            this->processDataLenght = this->tailDataLenght;
            CopyInTail(i, this->loopCount);
            Compute(i);
            CopyOut(i, this->loopCount);
        }
    }
private:
    __aicore__ inline void CopyIn(int32_t progress, int32_t step) {
        LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();
        LocalTensor<DTYPE_X> vyLocal = inQueueVy.AllocTensor<DTYPE_X>();
        LocalTensor<DTYPE_X> vxLocal = inQueueVx.AllocTensor<DTYPE_X>();
        DataCopy(xLocal, xGm[progress * this->interval * 2 + step*this->tileDataLenght], this->processDataLenght);
        DataCopy(xLocal[this->tileDataLenght], xGm[progress * this->interval * 2 + this->interval + step*this->tileDataLenght], this->processDataLenght);
        DataCopy(vyLocal[this->tileDataLenght], xGradGm[progress * this->interval * 2 + step*this->tileDataLenght], this->processDataLenght);
        DataCopy(vyLocal, vyGm[progress * this->interval + step*this->tileDataLenght], this->processDataLenght);
        DataCopy(vxLocal, vxGm[progress * this->interval * 2 + step*this->tileDataLenght], this->processDataLenght);
        DataCopy(vxLocal[this->tileDataLenght], vxGm[progress * this->interval * 2 + this->interval + step*this->tileDataLenght], this->processDataLenght);
        inQueueX.EnQue(xLocal);
        inQueueVy.EnQue(vyLocal);
        inQueueVx.EnQue(vxLocal);
    }
    __aicore__ inline void CopyInTail(int32_t progress, int32_t step) {
        LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();
        LocalTensor<DTYPE_X> vyLocal = inQueueVy.AllocTensor<DTYPE_X>();
        LocalTensor<DTYPE_X> vxLocal = inQueueVx.AllocTensor<DTYPE_X>();
        DataCopy(xLocal, xGm[progress * this->interval * 2 + step*this->tileDataLenght], this->tailDataLenght_32);
        DataCopy(xLocal[this->tileDataLenght], xGm[progress * this->interval * 2 + this->interval + step*this->tileDataLenght], this->tailDataLenght_32);
        DataCopy(vyLocal[this->tileDataLenght], xGradGm[progress * this->interval * 2 + step*this->tileDataLenght], this->tailDataLenght_32);
        DataCopy(vyLocal, vyGm[progress * this->interval + step*this->tileDataLenght], this->tailDataLenght_32);
        DataCopy(vxLocal, vxGm[progress * this->interval * 2 + step*this->tileDataLenght], this->tailDataLenght_32);
        DataCopy(vxLocal[this->tileDataLenght], vxGm[progress * this->interval * 2 + this->interval + step*this->tileDataLenght], this->tailDataLenght_32);
        inQueueX.EnQue(xLocal);
        inQueueVy.EnQue(vyLocal);
        inQueueVx.EnQue(vxLocal);
    }
    __aicore__ inline void Compute(int32_t progress) {
        LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>();
        LocalTensor<DTYPE_X> vyLocal = inQueueVy.DeQue<DTYPE_X>();
        LocalTensor<DTYPE_X> vxLocal = inQueueVx.DeQue<DTYPE_X>();
        LocalTensor<DTYPE_X> outLocal = outQueueOut.AllocTensor<DTYPE_X>();
        auto tmp1 = QueueTmp1.Get<float>();
        auto tmp2 = QueueTmp2.Get<float>();
        auto tmp3 = QueueTmp3.Get<float>();
        Cast(tmp1, xLocal[this->tileDataLenght], RoundMode::CAST_NONE, this->processDataLenght);
        Muls(tmp1, tmp1, (float)(-1), this->processDataLenght);
        Exp(tmp1, tmp1, this->processDataLenght);
        Adds(tmp1, tmp1, (float)(1), this->processDataLenght);
        Duplicate(tmp2, (float)(1), this->processDataLenght);
        Div(tmp1, tmp2, tmp1, this->processDataLenght);
        Cast(tmp3, vxLocal[this->tileDataLenght], RoundMode::CAST_NONE, this->processDataLenght);
        Mul(tmp2, tmp3, tmp1, this->processDataLenght);
        Sub(tmp2, tmp3, tmp2, this->processDataLenght);
        Cast(outLocal.ReinterpretCast<float>(), vyLocal, RoundMode::CAST_NONE, this->processDataLenght);
        Mul(outLocal.ReinterpretCast<float>(), outLocal.ReinterpretCast<float>(), tmp1, this->processDataLenght);
        Cast(tmp3, vyLocal[this->tileDataLenght], RoundMode::CAST_NONE, this->processDataLenght);
        Mul(vyLocal.ReinterpretCast<float>(), tmp3, tmp2, this->processDataLenght);
        Add(vyLocal.ReinterpretCast<float>(), outLocal.ReinterpretCast<float>(), vyLocal.ReinterpretCast<float>(), this->processDataLenght);
        Cast(outLocal.ReinterpretCast<float>(), xLocal, RoundMode::CAST_NONE, this->processDataLenght);
        Mul(xLocal.ReinterpretCast<float>(), tmp1, outLocal.ReinterpretCast<float>(), this->processDataLenght);
        Sub(outLocal.ReinterpretCast<float>(), outLocal.ReinterpretCast<float>(), xLocal.ReinterpretCast<float>(), this->processDataLenght);
        Mul(tmp2, tmp2, xLocal.ReinterpretCast<float>(), this->processDataLenght);
        Cast(xLocal.ReinterpretCast<float>(), vxLocal, RoundMode::CAST_NONE, this->processDataLenght);
        Mul(tmp1, tmp1, xLocal.ReinterpretCast<float>(), this->processDataLenght);
        Sub(tmp1, xLocal.ReinterpretCast<float>(), tmp1, this->processDataLenght);
        Sub(tmp1, tmp1, tmp2, this->processDataLenght);
        Mul(tmp1, tmp1, tmp3, this->processDataLenght);
        Mul(tmp2, vyLocal.ReinterpretCast<float>(), outLocal.ReinterpretCast<float>(), this->processDataLenght);
        Add(tmp1, tmp1, tmp2, this->processDataLenght);
        Cast(outLocal, vyLocal.ReinterpretCast<float>(), RoundMode::CAST_RINT, this->processDataLenght);
        Cast(outLocal[this->tileDataLenght], tmp1, RoundMode::CAST_RINT, this->processDataLenght);
        inQueueX.FreeTensor(xLocal);
        inQueueVy.FreeTensor(vyLocal);
        inQueueVx.FreeTensor(vxLocal);
        outQueueOut.EnQue<DTYPE_X>(outLocal);
    }
    __aicore__ inline void CopyOut(int32_t progress, int32_t step) {
        LocalTensor<DTYPE_X> outLocal = outQueueOut.DeQue<DTYPE_X>();
        DataCopyExtParams copyParams{1, static_cast<uint32_t>(this->processDataLenght * sizeof(DTYPE_X)), 0, 0, 0}; 
        DataCopyPad(outGm[progress * this->interval * 2 + step*this->tileDataLenght], outLocal, copyParams);
        DataCopyPad(outGm[progress * this->interval * 2 + this->interval + step*this->tileDataLenght], outLocal[this->tileDataLenght], copyParams);
        outQueueOut.FreeTensor(outLocal);
    }
private:
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueVy, inQueueVx;
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueOut;
    TBuf<QuePosition::VECCALC> QueueTmp1, QueueTmp2,QueueTmp3,QueueTmp4, QueueTmp5;
    GlobalTensor<DTYPE_X> xGradGm;
    GlobalTensor<DTYPE_X> xGm;
    GlobalTensor<DTYPE_X> vyGm;
    GlobalTensor<DTYPE_X> vxGm;
    GlobalTensor<DTYPE_X> outGm;
    uint32_t repeat;
    uint32_t interval;
    uint32_t loopCount;
    uint32_t tileDataLenght;
    uint32_t tailDataLenght;
    uint32_t tailDataLenght_32;
    uint32_t processDataLenght;
};
extern "C" __global__ __aicore__ void glu_grad_jvp(GM_ADDR x_grad, GM_ADDR y_grad, GM_ADDR x, GM_ADDR v_y, GM_ADDR v_x, GM_ADDR jvp_out, GM_ADDR workspace, GM_ADDR tiling) {
    GET_TILING_DATA(tiling_data, tiling);
    TPipe pipe;
    if(TILING_KEY_IS(1)) { 
        KernelGluGradJvp_float1 op;
        op.Init(x_grad, y_grad, x, v_y, v_x, jvp_out,
                tiling_data.bigNum, tiling_data.bigLength, tiling_data.smallLength, 
                tiling_data.interval, &pipe);
        op.Process();
    }
    else if(TILING_KEY_IS(2)) { 
        KernelGluGradJvp_float op;
        op.Init(x_grad, y_grad, x, v_y, v_x, jvp_out,
                tiling_data.bigNum, tiling_data.bigLength, tiling_data.smallLength, 
                tiling_data.interval, &pipe);
        op.Process();
    }
    else if(TILING_KEY_IS(3)) { 
        KernelGluGradJvp_half op;
        op.Init(x_grad, y_grad, x, v_y, v_x, jvp_out,
                tiling_data.bigNum, tiling_data.bigLength, tiling_data.smallLength, 
                tiling_data.interval, &pipe);
        op.Process();
    }
}